/******************************************************************************
*  Copyright(c) 2021-2024 Sietium
*
*  @file:    gmem-read-bw.cpp
*  @brif:    内存读带宽
*
*  @date:    09/22/2022
*  @Author:  way.lau
*
******************************************************************************/
#include <stdlib.h>
#include <string.h>
#include "objectfactory.h"
#include "node.h"
#include "clcheck.h"
#include "cldeviceinfo.h"
#include "clevent.h"
#include "cltkit.h"

#define REPART 50

class GMem_read_bw : public Node
{
public:
    GMem_read_bw()
    {
        nodeUint = "GBPS";
    }
    std::string name() override
    {
        return "gmem-rw-bandwidth";
    }
    std::string describe() override
    {
        return "内存读写带宽";
    }
    Node::NodeType  nodeType() override
    {
        return NodeTP_MemRw;
    }
protected:
    bool startup() override;
    bool perform() override;
    void shutdown() override;
private:
    std::string kernelCode();
private:
    struct cl_handle_t {
        cl_device_id        deviceid;
        cl_context          context;
        cl_command_queue    queue;
    } handle;
};

bool GMem_read_bw::startup()
{
    memset(&handle, 0, sizeof(cl_handle_t));

    handle.deviceid = deviceinfo->deviceID();
    handle.context = clCreateContext(NULL, 1,  &handle.deviceid, NULL, NULL, NULL);
    handle.queue = clCreateCommandQueue(handle.context, handle.deviceid,
                                        CL_QUEUE_PROFILING_ENABLE,
                                        NULL);
    return true;
}

bool GMem_read_bw::perform()
{
    cl_int err;
    cl_ulong globalWIs;
    cl_ulong localWIs;
    CLTkit::getNDRangeWorkItemSize(deviceinfo->maxComputeUnit(),
                                   deviceinfo->maxWGSize(),
                                   deviceinfo->maxAllocSize(),
                                   deviceinfo->deviceType(),
                                   sizeof(cl_float),
                                   globalWIs,
                                   localWIs);
    //create and init mem
    char *buff = new  char[globalWIs];
    memset(buff, 0x12, globalWIs);
    clFlush(handle.queue);
    cl_mem meminput   = clCreateBuffer(handle.context, CL_MEM_READ_WRITE,
                                       sizeof(cl_float) * globalWIs, NULL, &err);

    CHECK_CL_ERROR(err);
    CHECK_CL_ERROR(clEnqueueWriteBuffer(handle.queue, meminput, CL_TRUE, 0, globalWIs,
                                        buff, 0,
                                        NULL,
                                        NULL));
    cl_mem memoutput   = clCreateBuffer(handle.context, CL_MEM_WRITE_ONLY,
                                        sizeof(cl_float) * globalWIs, NULL, &err);
    CHECK_CL_ERROR(err);
    // Create program object
    std::string code = kernelCode();
    const char *progSource[] = {code.c_str()};
    cl_program program = clCreateProgramWithSource(handle.context, 1, progSource,
                                                   NULL, &err);
    CHECK_CL_ERROR(err);
    // Compile the program
    CHECK_CL_ERROR(clBuildProgram(program, 1, &handle.deviceid, NULL, NULL, NULL));

    //create kernel
    cl_kernel kernel = clCreateKernel(program, "gmemrwtest", NULL);
    //set kernel arg
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &meminput);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &memoutput);

    cl_ulong ltime = 0.0;
    for (int i = 0; i < REPART; ++i) {
        size_t global_work_size[1] = {globalWIs};
        size_t local_work_size[1]  = {localWIs};
        CLEvent evt;
        err = clEnqueueNDRangeKernel(handle.queue, kernel, 1, NULL, global_work_size, local_work_size, 0,
                                     NULL,
                                     evt.event());
        CHECK_CL_ERROR(err);
        CHECK_CL_ERROR(clFinish(handle.queue));
        if (evt.profiling())
            ltime += (evt.endTime() - evt.submitTime());
    }

    double fbw =  (globalWIs * sizeof(float) * 2) / (ltime / REPART / 1e9f);
    nodeValue = fbw / 1e9f;
    // release obj
    delete []buff;
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseMemObject(meminput);
    clReleaseMemObject(memoutput);
    return true;
}

void GMem_read_bw::shutdown()
{
    if (handle.queue) {
        CHECK_CL_ERROR(clReleaseCommandQueue(handle.queue));
        handle.queue = nullptr;
    }
    if (handle.context) {
        CHECK_CL_ERROR(clReleaseContext(handle.context));
        handle.context = nullptr;
    }
}

std::string GMem_read_bw::kernelCode()
{
    const char *kernelSource =
        "__kernel void gmemrwtest(__global const float* input,__global float* output)"
        "{                                                            "
        "    int idx = get_global_id(0);                              "
        "    float sum = input[idx] + 1.2;                            "
        "    output[idx] = sum;                                       "
        "} ";
    return std::string(kernelSource);
}

REGISTER_OBJECT(Node, GMem_read_bw)
